#ifndef M_PI
#define M_PI 0x1.921fb54442d18p+1
#endif
#define complexMul(a,b) ((float2)(mad(-(a).y, (b).y, (a).x * (b).x), mad((a).y, (b).x, (a).x * (b).y)))
                          
__kernel void ReorderElements(__global float2* x, int num_bits) {
  int idx = get_global_id(0);

  int r = 0;
  int k = idx;
  for (int i = 0; i < num_bits; ++i) {
    r <<= 1;
    r |= k & 1;
    k >>= 1;
  }

  if (idx < r) {
    float2 temp = x[idx];
    x[idx] = x[r];
    x[r] = temp;
  }
}


__kernel void DftOnePass(__global float2* x, int part_sz_log) {
  int idx = get_global_id(0);
  int i = idx >> part_sz_log;
  int j = idx - (i << part_sz_log);
  i <<= part_sz_log + 1;

  int sz = 1 << part_sz_log;
  float2 a = x[i + j];
  float2 b = x[i + j + sz];

  float ang = -(float)M_PI / sz * j;
  b = complexMul(b, (float2)(native_cos(ang), native_sin(ang)));

  x[i + j] = a + b;
  x[i + j + sz] = a - b;
}
